

//#pragma OPENCL EXTENSION cl_amd_printf : enable


#include "cl_config.h"
#include "geometry.h"
#include "shape_funcs.h"
#include "primitive_funcs.h"
#include "ray.h"

__kernel void ray_intersect(
		GLOBAL intersection_t *intersections,
		GLOBAL ray_t *rays, 
		GLOBAL const float* shape_data,
		GLOBAL const float* accelerator_data,
		GLOBAL const primitive_info_t* primitives,
		const unsigned int primitive_count, 
		const unsigned int number_work_items,
		LOCAL int *local_stack
)		
{
	const int gid = get_global_id(0);
	if (gid >= number_work_items) return;
	ray_t ray;
	rassign(ray,rays[gid]);
	ray.d.w = 0;
	ray.o.w = 0;
//	intersection_t isect;
//	isect.primitive_idx = 0xffffffff;
	intersect(accelerator_data, shape_data, primitives, primitive_count,&ray,intersections + gid, local_stack);
//	intersections[gid] = isect;
//	intersections[gid].primitive_idx = isect.primitive_idx;
}


